#include <stdio.h>
#include <stdlib.h>
#include <sys/types.h>
#include <sys/stat.h>
#include <fcntl.h>
#include <sys/mman.h>
#include <string.h>
#include <unistd.h>
#include <time.h>
#include "scanImpl.cu"

#define LTH	1
#define	GEQ	2
#define BILLION 1000000000


#  define CUDA_SAFE_CALL_NO_SYNC( call) do {                                \
    cudaError err = call;                                                    \
    if( cudaSuccess != err) {                                                \
        fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n",        \
                __FILE__, __LINE__, cudaGetErrorString( err) );              \
        exit(EXIT_FAILURE);                                                  \
    } } while (0)
#  define CUDA_SAFE_CALL( call) do {                                        \
    CUDA_SAFE_CALL_NO_SYNC(call);                                            \
    cudaError err = cudaThreadSynchronize();                                 \
    if( cudaSuccess != err) {                                                \
        fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n",        \
                __FILE__, __LINE__, cudaGetErrorString( err) );              \
        exit(EXIT_FAILURE);                                                  \
    } } while (0)


__global__ static void genScanFilter(char *col,  int tupleNum, int rate, int * filter, int *count, int * totalCount){
        int stride = blockDim.x * gridDim.x;
        int tid = blockIdx.x * blockDim.x + threadIdx.x;
	int con;
	int min = 0;
	int max = rate;
	int localCount = 0;

	for(int i = tid; i<tupleNum;i+=stride){
		int a1 = ((int *)col)[i];
		con = (a1>=min);
		con &= (a1<max);
		filter[i] = con;
		if(con == 1)
			localCount ++;
	}

	count[tid] = localCount;
	atomicAdd(totalCount, localCount);
}


__global__ static void scan(char *col, int colSize, long tupleNum, int *psum, long resultNum, int * filter, char * result){
        int stride = blockDim.x * gridDim.x;
        int tid = blockIdx.x * blockDim.x + threadIdx.x; 
	int localCount = psum[tid];

	for(long i = tid; i<tupleNum;i+=stride){
		
		if(filter[i] == 1){
			((int *)result)[localCount] = ((int*)col)[i];
			localCount ++;
		}
	}
}

//needs to add support for compression later 

#define	RATE	0.01

int main(int argc, char **argv){

	struct timespec start,end;
	char ** column;
	int *gpuCount;
	int * gpuFilter;
	int  *gpuPsum;


	if(argc != 4){
		printf("./scanTest rate scanColNum whereColNum\n");
		exit(-1);
	}

	int *gpuTmp;

        CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **)&gpuTmp, sizeof(int)));
        CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpuTmp));

	int rate = atoi(argv[1]) * RATE * 100;
	int scanColNum = atoi(argv[2]);
	int whereColNum = atoi(argv[3]);

	int * scanCol = (int *) malloc(sizeof(int) * scanColNum);
	int * whereCol = (int *) malloc(sizeof(int) * whereColNum);

	if(!scanCol || !whereCol){
		printf("Failed to allocate host memory\n");
		exit(-1);
	}

	char **scanTable = (char **) malloc(sizeof(char *) * scanColNum);
	char **whereTable = (char **) malloc(sizeof(char *) * whereColNum);

	if(!scanTable || !whereTable){
		printf("Failed to allocate host memory\n");
		exit(-1);
	}

	for(int i=0;i<scanColNum;i++){
		char buf[32] = {0};
		sprintf(buf,"scanTable%d",i);
		scanCol[i] = open(buf,O_RDONLY);

		int tmpSize = lseek(scanCol[i],0,SEEK_END); 
		//scanTable[i] = (char*) malloc(tmpSize);
		cudaMallocHost((void**)&scanTable[i],tmpSize);
		if(!scanTable[i]){
			printf("Failed to allocate host memory for scanTable\n");
			exit(-1);
		}
		char *tmp = (char *)mmap(0,tmpSize,PROT_READ,MAP_SHARED,scanCol[i],0);
		memcpy(scanTable[i], tmp, tmpSize);
		munmap(tmp,tmpSize);
	}

	for(int i=0;i<whereColNum;i++){
		char buf[32] = {0};
		sprintf(buf,"whereTable%d",i);
		whereCol[i] = open(buf,O_RDONLY);

		int tmpSize = lseek(whereCol[i],0,SEEK_END); 
		//whereTable[i] = (char*) malloc(tmpSize);
		cudaMallocHost((void**)&whereTable[i],tmpSize);
		if(!whereTable[i]){
			printf("Failed to allocate host memory for whereTable\n");
			exit(-1);
		}
		char *tmp = (char *)mmap(0,tmpSize,PROT_READ,MAP_SHARED,whereCol[i],0);
		memcpy(whereTable[i], tmp, tmpSize);
		munmap(tmp,tmpSize);
	}

	clock_gettime(CLOCK_REALTIME,&start);
	dim3 grid(1024);
	dim3 block(256);

	int threadNum = grid.x * block.x;
	int whereSize = lseek(whereCol[0],0,SEEK_END);
	int totalTupleNum = whereSize / sizeof(int); 

	column = (char **) malloc(whereColNum * sizeof(char *));
	for(int i=0; i<whereColNum; i++){
		CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **) &column[i] , whereSize));
		CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(column[i], whereTable[i], whereSize, cudaMemcpyHostToDevice));
	}
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **)&gpuFilter,sizeof(int) * totalTupleNum));
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void**)&gpuPsum,sizeof(int)*threadNum));
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void**)&gpuCount,sizeof(int)*threadNum));
	CUDA_SAFE_CALL_NO_SYNC(cudaMemset(gpuPsum,0,sizeof(int) * threadNum));

	int count, *gpuTotalCount;
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **)&gpuTotalCount, sizeof(int)));
	CUDA_SAFE_CALL_NO_SYNC(cudaMemset(gpuTotalCount, 0 ,sizeof(int)));

	genScanFilter<<<grid,block>>>(column[0],totalTupleNum, rate, gpuFilter,gpuCount,gpuTotalCount);

	scanImpl(gpuCount,threadNum,gpuPsum);
	CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(&count, gpuTotalCount, sizeof(int), cudaMemcpyDeviceToHost));

	CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpuCount));
	CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpuTotalCount));

	printf("%d\n",count);

	for(int i=0; i<whereColNum;i++){
		CUDA_SAFE_CALL_NO_SYNC(cudaFree(column[i]));
	}

	free(column);

	char **result;

	column = (char **) malloc(scanColNum * sizeof(char *));
	result = (char **) malloc(scanColNum * sizeof(char *));

	for(int i=0;i<scanColNum;i++){
		CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **) &column[i] , whereSize));
		CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(column[i], scanTable[i], whereSize, cudaMemcpyHostToDevice));

		int colSize = count * sizeof(int); 

		CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **) &result[i], colSize));
		scan<<<grid,block>>>(column[i], sizeof(int), totalTupleNum, gpuPsum,count, gpuFilter,result[i]);

		CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(scanTable[i],result[i],colSize ,cudaMemcpyDeviceToHost));
		CUDA_SAFE_CALL(cudaFree(result[i]));
		CUDA_SAFE_CALL(cudaFree(column[i]));
	}


	CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpuPsum));
	CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpuFilter));

	free(column);
	free(result);


	for(int i=0;i<scanColNum;i++){
		close(scanCol[i]);
		cudaFreeHost(scanTable[i]);
	}

	for(int i=0;i<whereColNum;i++){
		close(whereCol[i]);
		cudaFreeHost(whereTable[i]);
	}
	free(scanTable);
	free(whereTable);

	clock_gettime(CLOCK_REALTIME,&end);
        double timeE = (end.tv_sec -  start.tv_sec)* BILLION + end.tv_nsec - start.tv_nsec;
        printf("Table scan CPU time: %lf\n", timeE);
}

